Skip to content

GPU dynamic memory allocation - #18

Merged
highperformancecoder merged 19 commits into
masterfrom
copilot-question
Aug 10, 2026
Merged

GPU dynamic memory allocation#18
highperformancecoder merged 19 commits into
masterfrom
copilot-question

Conversation

@highperformancecoder

@highperformancecoder highperformancecoder commented Aug 10, 2026

Copy link
Copy Markdown
Owner

This change is Reviewable

Summary by CodeRabbit

  • Performance and Reliability

    • Improved device memory allocation and recycling for more efficient accelerator workloads.
    • Added synchronization improvements to ensure parallel operations complete reliably.
    • Optimized array copying, resizing, swapping, and reference management.
  • API Improvements

    • Added local thread identification for parallel execution.
    • Simplified allocator behavior across host and device environments.
    • Updated mutation and population-count interfaces.
  • Simulation Updates

    • Increased spatial grid resolution to 8×8.
    • Simulations now process more items per step and report species counts.
    • Added per-step timing output for performance monitoring.

Copilot AI lite review requested due to automatic review settings August 10, 2026 23:21
@coderabbitai

coderabbitai Bot commented Aug 10, 2026

Copy link
Copy Markdown

Review Change Stack

Warning

Review limit reached

@highperformancecoder, you've reached your PR review limit, so we couldn't start this review.

Next review available in: 27 minutes

You've used all free OSS reviews for now. Wait for the free limit to reset to keep reviewing this public repository.

How can I continue?

After more reviews become available, a review can be triggered using the @coderabbitai review command as a PR comment. Alternatively, push new commits to this PR.

To avoid repeated limits, reduce automatic review volume by pausing incremental auto-reviews earlier, using label-based review opt-in, excluding WIP or generated PR titles, or requesting reviews manually when the PR is ready. If your team needs uninterrupted high-volume reviews, an organization admin can enable usage-based reviews.

How do review limits work?

CodeRabbit enforces per-developer PR review limits for each organization. Most developers receive the normal plan review availability.

For paid Pro and Pro+ PR reviews, CodeRabbit uses adaptive limits for sustained high-volume activity. When a developer's recent PR review activity reaches the 95th percentile or higher among CodeRabbit users, additional reviews become available more gradually as earlier reviews age out of the rolling window.

Please refer docs for additional details.

Review details
⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Pro Plus

Run ID: dfc71540-f921-43ea-8cde-1d4c7bef72d2

📥 Commits

Reviewing files that changed from the base of the PR and between 288c328 and ba79475.

📒 Files selected for processing (2)
  • include/arrays.h
  • include/sycl.h
📝 Walkthrough

Walkthrough

The change replaces the device allocation queue with active and discard stacks, adds SYCL synchronization and atomic reference handling for arrays, updates spatial mutation to use allocator-aware arrays, and changes the spatial benchmark configuration and timing output.

Changes

SYCL allocation and spatial model

Layer / File(s) Summary
Device allocator runtime
include/DeviceAllocator.h, include/non-sycl.h, include/sycl.h
Device allocation now uses LIFO active and discard stacks. Initialization uses device compute units. Discarded blocks are recycled. Host and device allocator aliases are separated.
SYCL array ownership and synchronization
include/arrays.h
Array reference counting, copying, resizing, swapping, and assignment now use SYCL atomics and work-group barriers.
SYCL execution and allocator recycling
include/ecolab.h
SYCL execution waits for completion. Work-group sizing uses device compute units. Fatal-error updates use device-scoped sequentially consistent atomics.
Allocator-aware spatial mutation
models/ecolab_model.h, models/ecolab_model.cc
Spatial counts and mutation results use UnsignedArray and density-derived allocators. Mutation uses scan buffers, group broadcasts, cleanup, and density checks.
Spatial benchmark execution
models/spatial_ecolab.py
The benchmark uses an 8×8 grid, generates 1000 items per step, records timing and species counts, and runs one step.

Classdesc revision

Layer / File(s) Summary
Classdesc submodule revision
classdesc
The classdesc submodule reference was updated to a different commit.

Estimated code review effort: 4 (Complex) | ~45 minutes

Sequence Diagram(s)

sequenceDiagram
  participant SpatialModel
  participant EcolabGraph
  participant SYCLQueue
  participant DeviceAllocator
  SpatialModel->>EcolabGraph: submit grouped mutation work
  EcolabGraph->>SYCLQueue: launch work groups
  EcolabGraph->>SYCLQueue: wait for completion
  EcolabGraph->>DeviceAllocator: recycle discard pile
Loading

Possibly related PRs

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Description Check ✅ Passed Check skipped - CodeRabbit’s high-level summary is enabled.
Title check ✅ Passed The title clearly describes the main changes to device allocation and SYCL memory management.
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
✨ Finishing Touches 💡 1
📝 Generate docstrings 💡
  • Create stacked PR
  • Commit on current branch
🧪 Generate unit tests (beta)
  • Create PR with unit tests
  • Commit unit tests in branch copilot-question

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands.

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

This PR advances the project’s GPU/SYCL support by refactoring the device-side allocator and related array semantics, and updates the spatial model to better operate in device-parallel contexts (plus adds timing instrumentation in the spatial Python driver).

Changes:

  • Reworked the SYCL device allocator to use a stack/discard-pile approach and added discard recycling after kernel execution.
  • Updated array ref-counting / barriers and SYCL helpers (e.g., localThreadId) to support device compilation paths.
  • Adjusted spatial model routines (nsp, mutate) for grouped execution and added per-step timing output in models/spatial_ecolab.py.

Reviewed changes

Copilot reviewed 9 out of 9 changed files in this pull request and generated 3 comments.

Show a summary per file
File Description
models/spatial_ecolab.py Changes grid sizing and adds timing + output logging for step execution.
models/ecolab_model.h Updates SYCL includes and adjusts method signatures for device-aware paths.
models/ecolab_model.cc Refactors nsp() and mutate() implementations for grouped/device execution.
include/sycl.h Replaces onDevice() with localThreadId() helper for device code.
include/non-sycl.h Adjusts non-SYCL stubs/aliases (moves LocalAllocator alias here).
include/ecolab.h Adds extra SYCL synchronization and allocator discard recycling after kernels.
include/DeviceAllocator.h Major allocator refactor: queue→stack, discard pile, allocator construction changes.
include/arrays.h Refactors COW ref-count access to be SYCL-safe (atomic/barrier changes).
Suppressed comments (1)

include/DeviceAllocator.h:70

  • Stack::pop uses t++ and assigns t=size, but sycl::atomic_ref doesn’t support these operators/assignments. This won’t compile and also can’t safely restore top on underflow. Use fetch_add/fetch_sub with a bounds check instead.
    unsigned pop()
    {
      Atomic t(top);
      unsigned p=t++;
      if (p>=size) {t=size; return ~0;} // stack empty
      return slots[p];
    }

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Comment thread include/arrays.h
Comment thread models/spatial_ecolab.py
Comment on lines +59 to +68
from time import perf_counter

epoch=2000000
mut_factor=1000

extinctions=0
migrations=0

out=open("time.dat","w")

Comment thread include/DeviceAllocator.h

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 9

🧹 Nitpick comments (3)
include/ecolab.h (1)

156-156: 📐 Maintainability & Code Quality | 🔵 Trivial

Track the missing const overload.

This TODO leaves read-only graph traversal unavailable through forAll. Add a const-safe overload or track this work in an issue.

Do you want me to prepare the overload and its tests?

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@include/ecolab.h` at line 156, Implement the missing const-safe overload for
forAll using a free template function that preserves constness of the graph’s
this pointer and passes it through to the functional’s arguments; add coverage
for read-only graph traversal, or document the work in an issue if the overload
cannot be implemented now.
include/arrays.h (1)

1712-1715: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Remove the obsolete condition comment. onDevice() has no remaining declarations, definitions, or callers. Use dt=x.dt; directly.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@include/arrays.h` around lines 1712 - 1715, In the surrounding code, remove
the obsolete commented-out condition and retain the unconditional assignment
`dt=x.dt;` directly before `incrRef()`.
include/non-sycl.h (1)

26-33: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Delete the dead GroupLocal block. The current include graph keeps the LocalAllocator aliases separate. DeviceAllocator.h is included only when SYCL_LANGUAGE_VERSION is defined, while non-sycl.h is included only in the non-SYCL branch.

No alias-collision fix is required for the current include graph.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@include/non-sycl.h` around lines 26 - 33, Remove the commented-out GroupLocal
template block immediately before SyclRandomEngine in non-sycl.h. Leave
SyclRandomEngine and LocalAllocator unchanged; no alias-collision changes are
needed.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@include/arrays.h`:
- Around line 1558-1572: Update both affected sites in
include/arrays.h:1558-1572 and include/arrays.h:1653-1668. In the array cleanup
path around ref() and decrRef(), reduce the dt-null condition across the group
before branching, move the barrier after the branch so it is unconditional, and
preserve the existing reference-count cleanup. In the second path, replace the
early return with a group-uniform condition and ensure the barrier executes
unconditionally so every work item reaches the same number of barriers.
- Around line 1595-1601: Update copy() to check the result of alloc(sz) before
dereferencing it, matching the existing null guard in resize(). If allocation
fails, preserve the allocator’s fatal-error behavior and return without
accessing dt->dt; only perform the SYCL assignment or memcpy when the allocation
succeeds.

In `@include/DeviceAllocator.h`:
- Line 200: Update the comment on LocalAllocatorSize to accurately state that
the value is 8 KiB, replacing the incorrect 32 KiB description while preserving
the constant and its occupancy-related usage.
- Around line 38-70: Update the Stack comment to describe a bounded LIFO stack,
its push/pop operations, and the constraint that concurrent device pushes and
pops must not target the same stack. In Stack::push, atomically validate that
top is greater than zero before decrementing and writing; reject or otherwise
safely handle underflow without indexing slots. Preserve the existing full-stack
behavior and avoid publishing an invalid slot.
- Around line 171-178: Update GlobalDeviceAllocator’s __SYCL_DEVICE_ONLY__
construction path so it cannot leave allocator as nullptr: delete the device
default constructor or make allocation fail in a controlled manner before
dereferencing allocator. Ensure allocate() handles any unavailable device
allocator safely while preserving the existing host initialization through
deviceAllocator().
- Around line 52-57: Update DeviceAllocator’s Stack::init to accept the
simple-range kernel’s sycl::id<1> or sycl::item<1> argument instead of calling
syclItem(). Derive the starting index and stride from that argument and the
kernel launch size, while preserving the existing slot initialization and top
reset behavior.

In `@include/sycl.h`:
- Around line 141-144: The device branch of groupLeader() must return the SYCL
group’s actual leader status instead of falling through to unconditional true;
restore an explicit device return such as syclGroup().leader() so exactly one
work item is selected. In include/sycl.h lines 141-144, update groupLeader(); in
include/DeviceAllocator.h lines 137-145, verify the discard push remains guarded
by groupLeader() and requires no direct change if corrected by that root fix.

In `@models/ecolab_model.h`:
- Line 145: Restore const access for SpatialModel::nsp by adding a
const-qualified overload while preserving the existing non-const overload and
behavior. Ensure callers with a const SpatialModel can invoke nsp without
qualification errors.

In `@models/spatial_ecolab.py`:
- Line 67: Make the time.dat handling in the timing flow rank-safe: guard the
open operation and writes around out and the lines 84-85 timing output with
myid()==0, so only rank 0 creates and writes the shared file. Preserve the
existing timing behavior while preventing other ranks from truncating or
overwriting it.

---

Nitpick comments:
In `@include/arrays.h`:
- Around line 1712-1715: In the surrounding code, remove the obsolete
commented-out condition and retain the unconditional assignment `dt=x.dt;`
directly before `incrRef()`.

In `@include/ecolab.h`:
- Line 156: Implement the missing const-safe overload for forAll using a free
template function that preserves constness of the graph’s this pointer and
passes it through to the functional’s arguments; add coverage for read-only
graph traversal, or document the work in an issue if the overload cannot be
implemented now.

In `@include/non-sycl.h`:
- Around line 26-33: Remove the commented-out GroupLocal template block
immediately before SyclRandomEngine in non-sycl.h. Leave SyclRandomEngine and
LocalAllocator unchanged; no alias-collision changes are needed.
🪄 Autofix

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Pro Plus

Run ID: 596c5d58-bc1f-49ef-be56-766edc52e592

📥 Commits

Reviewing files that changed from the base of the PR and between 0d597ce and 288c328.

📒 Files selected for processing (9)
  • classdesc
  • include/DeviceAllocator.h
  • include/arrays.h
  • include/ecolab.h
  • include/non-sycl.h
  • include/sycl.h
  • models/ecolab_model.cc
  • models/ecolab_model.h
  • models/spatial_ecolab.py

Comment thread include/arrays.h
Comment on lines +1558 to +1572
groupBarrier();
if (dt)
{
if (ref()==1)
unsigned refCnt=ref();
#ifdef __SYCL_DEVICE_ONLY__
// ensure absolute consistency of the reference count
refCnt=sycl::reduce_over_group(syclGroup(),refCnt,sycl::minimum<unsigned>());
#endif
if (refCnt==1)
{
free(dt);
return;
}
decrRef();
dt=nullptr;
} else
decrRef();
groupBarrier();

@coderabbitai coderabbitai Bot Aug 10, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🩺 Stability & Availability | 🔴 Critical | 🏗️ Heavy lift

Work-group barriers run inside data-dependent branches in array. Both sites call groupBarrier() from control flow that is not proven uniform across the work group. swap() states that an array may be per work item or shared by the group, so dt, size() and ref() can differ between work items. A barrier that only some work items reach is undefined behaviour and normally hangs the kernel. Hoist every barrier to unconditional control flow, or reduce the branch condition over the group first so all work items take the same path.

  • include/arrays.h#L1558-L1572: move the barrier on line 1572 out of the if (dt) block, and reduce the null test of dt over the group before branching.
  • include/arrays.h#L1653-L1668: replace the early return on line 1654 with a group-uniform condition, and make the barrier on line 1661 unconditional so both paths execute the same barrier count.
📍 Affects 1 file
  • include/arrays.h#L1558-L1572 (this comment)
  • include/arrays.h#L1653-L1668
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@include/arrays.h` around lines 1558 - 1572, Update both affected sites in
include/arrays.h:1558-1572 and include/arrays.h:1653-1668. In the array cleanup
path around ref() and decrRef(), reduce the dt-null condition across the group
before branching, move the barrier after the branch so it is unconditional, and
preserve the existing reference-count cleanup. In the second path, replace the
early return with a group-uniform condition and ensure the barrier executes
unconditionally so every work item reaches the same number of barriers.

Copy link
Copy Markdown
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The only way for dt to be null is for the array to be created that way, for alloc to be called with zero as an argument or for line 1569 to be hit. In the first two cases, all threads participate. In the last case, the reduce_over_group is an implied barrier, so all threads must have reached that point before line 1569 is hit. In which case, all threads will reach the groupBarrier on line 1572.

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@highperformancecoder, agreed for array::release().

With the stated dt lifecycle invariant, if (dt) is work-group uniform. In the free path, all work items must participate in sycl::reduce_over_group before any work item can set dt to nullptr. Therefore, all work items reach the barrier after the branch.

I withdraw the include/arrays.h release() portion of the finding. The separate resize() path still needs its own uniformity argument.

You are interacting with an AI system.

Comment thread include/arrays.h
Comment thread include/DeviceAllocator.h
Comment thread include/DeviceAllocator.h
Comment on lines 52 to 57
void init() {
top=0; // full stack
for (unsigned i=syclItem().get_global_linear_id(); i<size;
i+=syclItem().get_global_range().size()) {
slots[i].value=i;
slots[i].seq=i+1;
}
i+=syclItem().get_global_range().size())
slots[i]=i;
}

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🩺 Stability & Availability | 🔴 Critical | ⚡ Quick win

🧩 Analysis chain

🌐 Web query:

SYCL sycl::ext::oneapi::this_work_item::get_nd_item requirements nd_range kernel simple range

💡 Result:

The function sycl::ext::oneapi::this_work_item::get_nd_item is a SYCL extension used to retrieve a sycl::nd_item object representing the current work-item within a kernel [1][2]. Requirements and Usage: 1. Kernel Type: It is designed specifically for use within ND-range kernels [3][4][5]. It is not supported for simple range kernels (kernels launched with a simple sycl::range) [4]. 2. ND-Range Launch: The kernel must be launched with a sycl::nd_range [3]. 3. Dimensionality: The template parameter Dimensions must match the dimensionality of the currently executing kernel [3]. 4. Context: It is most commonly used within "free function" kernels (a SYCL extension) where the sycl::nd_item is not automatically passed as an argument to the function, providing a way to query work-item information like global IDs or ranges from inside the kernel body [6][4][7]. When used correctly, this function returns a sycl::nd_item instance, allowing access to standard nd_item methods such as get_global_id, get_global_linear_id, get_local_id, and get_global_range [8][1][7]. Attempting to use this function in a context that was not launched as an ND-range kernel, such as a simple range kernel, is not supported and may lead to runtime errors or undefined behavior [4][5].

Citations:


🏁 Script executed:

#!/bin/bash
set -eu
printf '%s\n' '--- candidate files ---'
git ls-files | grep -E '(^|/)(DeviceAllocator\.h|sycl\.h)$' || true
printf '%s\n' '--- relevant symbols and call sites ---'
rg -n -C 6 'class Stack|struct Stack|void init|parallel_for|syclItem|get_nd_item' include/DeviceAllocator.h include/sycl.h 2>/dev/null || true
printf '%s\n' '--- file sizes ---'
wc -l include/DeviceAllocator.h include/sycl.h 2>/dev/null || true

Repository: highperformancecoder/ecolab

Length of output: 5694


🏁 Script executed:

#!/bin/bash
set -eu
printf '%s\n' '--- DeviceAllocator.h surrounding implementation ---'
cat -n include/DeviceAllocator.h | sed -n '1,150p'
printf '%s\n' '--- init call sites ---'
rg -n -C 5 '\.init\(\)|init\(\)' --glob '!include/DeviceAllocator.h' --glob '!include/sycl.h' .
printf '%s\n' '--- parallel_for signatures in repository ---'
rg -n -C 3 'parallel_for\s*\(' --glob '*.h' --glob '*.hpp' --glob '*.cpp' --glob '*.cc' --glob '*.cxx' .

Repository: highperformancecoder/ecolab

Length of output: 15838


🌐 Web query:

SYCL 2020 queue::parallel_for range overload kernel lambda argument sycl::item sycl::id simple range

💡 Result:

In SYCL 2020, the queue::parallel_for method allows you to enqueue a basic data-parallel kernel by specifying a execution range [1][2]. When using the overload that accepts a sycl::range, the kernel lambda function can accept either a sycl::id or a sycl::item as its argument [2]. sycl::id An id represents the unique index of the current work-item within the execution range [2]. It is a simple, lightweight way to access the current thread's coordinates [2][3]. Example using sycl::id: q.submit([&](sycl::handler &h) { h.parallel_for(sycl::range<1>(size), [=](sycl::id<1> idx) { // idx provides the current global index data_accessor[idx] += 5; }); }); sycl::item An item is a more feature-rich object that encapsulates both the work-item's index and the overall execution range [2][3]. It provides methods to query information about the execution space, such as the global range, without requiring additional parameters [4][5]. Example using sycl::item: q.submit([&](sycl::handler &h) { h.parallel_for(sycl::range<1>(size), [=](sycl::item<1> itm) { auto globalId = itm.get_id; // Equivalent to sycl::id auto range = itm.get_range; // Access to the range dimensions data_accessor[globalId] += 5; }); }); Key Considerations for SYCL 2020: - Offsets: Support for passing an offset to parallel_for has been deprecated in SYCL 2020 [6]. Developers are expected to perform any necessary offset arithmetic manually within the kernel lambda [6]. - Captures: When using a lambda function as a kernel, all captures should be by copy ([=]), and the lambda must not be mutable [7]. - Dimension constraints: SYCL typically supports up to 3 dimensions. For higher-dimensional needs, developers must use a 1D range and perform manual index linearization/calculation [6].

Citations:


🌐 Web query:

site:github.khronos.org/SYCL-Reference queue parallel_for range id item kernel function

💡 Result:

Relevant Khronos SYCL Reference pages:

  • sycl::queue::parallel_for — submits a kernel over a range or nd_range; returns an event. (github.khronos.org)
  • sycl::handler::parallel_for — shows kernels receiving sycl::item, with examples using range and id. (github.khronos.org)
  • sycl::item — provides each work-item’s ID and range via get_id() and get_range(). (github.khronos.org)
  • sycl::id — represents a position within a SYCL range. (github.khronos.org)
  • Defining kernels — explains lambda and function-object kernel functions. (github.khronos.org)

Basic pattern:

sycl::queue q;

q.submit([&](sycl::handler& h) {
    h.parallel_for<class Kernel>(
        sycl::range<1>(N),
        [=](sycl::id<1> i) {
            // kernel body; i[0] is the work-item index
        });
});

With sycl::item:

[=](sycl::item<1> item) {
    auto i = item.get_id()[0];
}

Citations:


Pass the simple-range index to Stack::init

syclItem() calls get_nd_item<1>(), which is not supported by the simple-range kernel at lines 112-113. Initialize the slots with the kernel’s sycl::id<1> or sycl::item<1> argument and use the launch size as the stride. Otherwise, allocator initialization can fail at runtime.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@include/DeviceAllocator.h` around lines 52 - 57, Update DeviceAllocator’s
Stack::init to accept the simple-range kernel’s sycl::id<1> or sycl::item<1>
argument instead of calling syclItem(). Derive the starting index and stride
from that argument and the kernel launch size, while preserving the existing
slot initialization and top reset behavior.

Comment thread include/DeviceAllocator.h
Comment on lines +171 to +178
DeviceAllocator<>* allocator;

GlobalDeviceAllocator() = default; // note: default constructor must be called on host
#ifdef __SYCL_DEVICE_ONLY__
GlobalDeviceAllocator(): allocator(nullptr) {} // = delete;
#else
GlobalDeviceAllocator() // note: default constructor must be called on host
{allocator=&deviceAllocator();}
#endif

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
# Find device-side default construction of arrays that use GlobalDeviceAllocator.
rg -nP -C4 '\barray\s*<[^>]*(GlobalDeviceAllocator|Allocator<)[^>]*>\s+\w+\s*(;|\{\s*\})' --type=cpp --type=cc -g '!**/build/**'
# Inspect the terminating allocator and allocate() paths for null handling.
fd -t f 'DeviceAllocator.h' | xargs -r sed -n '80,200p'

Repository: highperformancecoder/ecolab

Length of output: 4390


🏁 Script executed:

#!/bin/bash
set -eu

echo '--- GlobalDeviceAllocator and array references ---'
rg -n -C3 'GlobalDeviceAllocator|m_allocator|array\s*\(' include src test examples --glob '!**/build/**' 2>/dev/null || true

echo '--- arrays.h structure and constructors ---'
fd -t f 'arrays.h' . | while read -r f; do
  echo "FILE: $f"
  ast-grep outline "$f" | head -120
  rg -n -C8 'm_allocator|array\s*\(|allocate\s*\(' "$f"
done

echo '--- array declarations and allocator aliases ---'
rg -n -C3 'using .*Allocator|typedef .*Allocator|array\s*<' . \
  --glob '*.{h,hpp,c,cpp,cc,cxx}' \
  --glob '!**/build/**' \
  --glob '!**/dist/**' 2>/dev/null || true

Repository: highperformancecoder/ecolab

Length of output: 50384


🏁 Script executed:

#!/bin/bash
set -eu

echo '--- All direct GlobalDeviceAllocator uses ---'
rg -n -C5 'GlobalDeviceAllocator' . --glob '!**/build/**' --glob '!**/dist/**'

echo '--- SYCL allocator aliases and device array declarations ---'
fd -t f -e h -e hpp -e cpp -e cc -e cxx . include src test models | while read -r f; do
  if rg -q 'DeviceAllocator|GlobalDeviceAllocator|__SYCL_DEVICE_ONLY__|array_ns::array|array<' "$f"; then
    echo "FILE: $f"
    rg -n -C4 'DeviceAllocator|GlobalDeviceAllocator|__SYCL_DEVICE_ONLY__|array_ns::array|array<' "$f" | head -160
  fi
done

echo '--- sycl.h relevant definitions ---'
fd -t f 'sycl.h' . | while read -r f; do
  echo "FILE: $f"
  sed -n '1,180p' "$f"
done

echo '--- arrays.h include and allocator-selection context ---'
sed -n '1,80p' include/arrays.h
sed -n '1240,1300p' include/arrays.h

Repository: highperformancecoder/ecolab

Length of output: 50383


🏁 Script executed:

#!/bin/bash
set -eu

echo '--- ModelData and EcolabPoint definitions ---'
sed -n '1,125p' models/ecolab_model.h

echo '--- DeviceType definition and uses ---'
rg -n -C8 'template.*DeviceType|struct DeviceType|class DeviceType|DeviceType<' include models src \
  --glob '*.{h,hpp,c,cpp,cc,cxx}' | head -240

echo '--- Device-side default construction of GlobalDeviceAllocator-backed objects ---'
rg -n -C8 'UnsignedArray|Allocator<|cell_ids|density|array<[^;]*(GlobalDeviceAllocator|ModelData::Allocator)' \
  models/ecolab_model.h models/ecolab_model.cc include \
  --glob '*.{h,hpp,c,cpp,cc,cxx}' | head -320

echo '--- Read-only semantic probe of the relevant preprocessor branches ---'
python3 - <<'PY'
from pathlib import Path

allocator = Path("include/DeviceAllocator.h").read_text()
arrays = Path("include/arrays.h").read_text()
model = Path("models/ecolab_model.h").read_text()

checks = {
    "device constructor initializes allocator to nullptr":
        "GlobalDeviceAllocator(): allocator(nullptr) {}" in allocator,
    "device allocate dereferences allocator without a guard":
        "{return reinterpret_cast<T*>(allocator->allocate(n*sizeof(T)));}" in allocator,
    "array default constructor is defaulted":
        "array()=default;" in arrays,
    "array stores allocator as a value member":
        "A m_allocator;" in arrays,
    "EcolabPoint uses GlobalDeviceAllocator under SYCL":
        "template <class T> using Allocator=GlobalDeviceAllocator<T>;" in model,
    "EcolabPoint has a GlobalDeviceAllocator-backed array member":
        "array<int,Allocator<int>> density;" in model,
}
for name, result in checks.items():
    print(f"{'PASS' if result else 'FAIL'}: {name}")
PY

Repository: highperformancecoder/ecolab

Length of output: 32452


Reject null device allocator construction.

When GlobalDeviceAllocator<T> is default-constructed on the device, allocator is nullptr, but allocate() dereferences it without a check. A default-constructed array<T, GlobalDeviceAllocator<T>> therefore fails on its first non-zero allocation. Delete the device constructor or report a controlled allocation failure.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@include/DeviceAllocator.h` around lines 171 - 178, Update
GlobalDeviceAllocator’s __SYCL_DEVICE_ONLY__ construction path so it cannot
leave allocator as nullptr: delete the device default constructor or make
allocation fail in a controlled manner before dereferencing allocator. Ensure
allocate() handles any unavailable device allocator safely while preserving the
existing host initialization through deviceAllocator().

Comment thread include/DeviceAllocator.h
};

constexpr static unsigned LocalAllocatorSize=30*1024; // 32KiB = half typical local storage
constexpr static unsigned LocalAllocatorSize=8*1024; // 32KiB = half typical local storage

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

The comment contradicts the value.

LocalAllocatorSize is now 8 KiB, but the comment states 32 KiB. include/ecolab.h derives wg_per_compute_unit from this constant, so a wrong comment misleads occupancy tuning.

📝 Proposed comment fix
-  constexpr static unsigned LocalAllocatorSize=8*1024; // 32KiB = half typical local storage
+  constexpr static unsigned LocalAllocatorSize=8*1024; // 8KiB per work group of local storage
📝 Committable suggestion

‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.

Suggested change
constexpr static unsigned LocalAllocatorSize=8*1024; // 32KiB = half typical local storage
constexpr static unsigned LocalAllocatorSize=8*1024; // 8KiB per work group of local storage
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@include/DeviceAllocator.h` at line 200, Update the comment on
LocalAllocatorSize to accurately state that the value is 8 KiB, replacing the
incorrect 32 KiB description while preserving the constant and its
occupancy-related usage.

Comment thread include/sycl.h Outdated
Comment thread models/ecolab_model.h
return *objects[makeId(x,y)];
}
array<unsigned> nsp() const;
array<unsigned> nsp();

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
set -euo pipefail

# Map nsp declarations and definitions before inspecting call sites.
ast-grep outline models --items all --type function --match nsp

# Inspect C++ call sites and nearby const-qualified model references.
rg -n --type=cpp -C4 '(?:\.|->)nsp\s*\(|const\s+SpatialModel' .

Repository: highperformancecoder/ecolab

Length of output: 710


🏁 Script executed:

#!/bin/bash
set -euo pipefail

printf '%s\n' '--- declarations and definitions ---'
rg -n -C3 '\bnsp\s*\(' models --glob '*.{h,hpp,hh,cc,cpp,cxx}' --glob '!*.md'

printf '%s\n' '--- all nsp call sites ---'
rg -n -C3 '\bnsp\s*\(' . --glob '*.{h,hpp,hh,cc,cpp,cxx,c}'

printf '%s\n' '--- SpatialModel type uses ---'
rg -n -C3 '\b(const\s+)?SpatialModel\b' . --glob '*.{h,hpp,hh,cc,cpp,cxx}'

Repository: highperformancecoder/ecolab

Length of output: 7604


Preserve const on SpatialModel::nsp. No in-repository callers use this method, but removing const breaks external callers with a const SpatialModel. Add a const overload.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@models/ecolab_model.h` at line 145, Restore const access for
SpatialModel::nsp by adding a const-qualified overload while preserving the
existing non-const overload and behavior. Ensure callers with a const
SpatialModel can invoke nsp without qualification errors.

Comment thread models/spatial_ecolab.py
@highperformancecoder
highperformancecoder merged commit 9383a1d into master Aug 10, 2026
2 of 3 checks passed
@highperformancecoder
highperformancecoder deleted the copilot-question branch August 10, 2026 23:53
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants